#ifdef __CUDACC__
extern "C"{
#endif

#include "common/common.h"
#include <cutil_inline.h>
#include <limits.h>
#include "mecu.h"

#define SEARCHWINDOW 8
#define LAMBDA 4
#define TILE 16

texture<uint8_t, 2, cudaReadModeElementType> TEXac;
texture<uint8_t, 2, cudaReadModeElementType> TEXan;

//---------------------------Declaracion de prototipos---------------------------

void EXcu(uint8_t *current,  uint8_t *ref, int strideCur, int strideRef, int width, int height,  int2** mvOut, int A,int2  *mvsX);
__global__ void SERCHcu(unsigned bw, unsigned bwUp, int2* inMV, int2* outMV);
__device__ unsigned MVCOSTcu(int dx, int dy);
__device__ unsigned SADcu(ImageTexture a, ImageTexture b, unsigned xStart, unsigned yStart, unsigned height, unsigned width, unsigned dx, unsigned dy);

//---------------------------Implementacion de Funciones---------------------------
__device__ unsigned MVCOSTcu(int dx, int dy){
	dx = (abs(dx))<<2;
	dy = (abs(dy))<<2;
	int xCost = round((log2f(dx+1)*2 + 0.718f + !!dx) + .5f);
	int yCost = round((log2f(dy+1)*2 + 0.718f + !!dy) + .5f);
	return (LAMBDA * (xCost+yCost));
}

__device__ unsigned SADcu(ImageTexture a, ImageTexture b, unsigned xStart, unsigned yStart, unsigned height, unsigned width, unsigned dx, unsigned dy){
	unsigned ret = 0;
	for(unsigned x = xStart; x< xStart + width; x++){
		for(unsigned y = yStart; y< yStart + height; y++){
			ret = __usad(tex2D(a, x,y),tex2D(b, x+dx,y+dy),ret);// |x-y| + z
		}
	}
	return ret;
}

//Metodo SERCH en cuda
__global__ void SERCHcu(unsigned bw, unsigned bwUp, int2* inMV, int2* outMV){
	const unsigned int bx = blockIdx.x;
	const unsigned int by = blockIdx.y;

	const unsigned int tx = threadIdx.x;
	const unsigned int ty = threadIdx.y;

	const int blockTop = TILE* by;//
	const int blockLeft = TILE* bx;

	const int dx = (tx-SEARCHWINDOW);//
	const int dy = (ty-SEARCHWINDOW);

	__shared__  unsigned mins[(2*SEARCHWINDOW+1)*(2*SEARCHWINDOW+1)];
	__shared__  int2 smallestLoc[(2*SEARCHWINDOW+1)*(2*SEARCHWINDOW+1)];

	unsigned myLoc =tx+ty *(2*SEARCHWINDOW+1);
	int2 oldMV = inMV[(bx>>1)+ (bwUp*(by>>1))];
	oldMV.x *= 2;
	oldMV.y *= 2;

	unsigned sum = SADcu(TEXac, TEXan, blockLeft,blockTop, TILE,TILE,dx+oldMV.x,dy+oldMV.y);
	
	mins[myLoc] = sum+ MVCOSTcu(dx,dy);
	
	smallestLoc[myLoc].x = dx+oldMV.x;
	smallestLoc[myLoc].y = dy+oldMV.y;

	__syncthreads();
	unsigned m=0;
	for(unsigned k = ((2*SEARCHWINDOW+1)*(2*SEARCHWINDOW+1));k>1;k=m){
		m = (k+1)>>1;
		if(myLoc<m && (m+myLoc<k))
		{
			if(mins[myLoc] > mins[myLoc+m]){
				smallestLoc[myLoc] = smallestLoc[myLoc+m];
				mins[myLoc] = mins[myLoc+m];
			}
		}

		__syncthreads();
	}

	if(myLoc==0){
		outMV[bx+by*bw] = smallestLoc[0];
	}
}
//Ejecuta en el kernel
void EXcu(uint8_t *current,  uint8_t *ref, int strideCur, int strideRef, int width, int height,  int2** retD, int A,int2  *inMVD){

    int bw = (width + (TILE<<A)-1)/((TILE<<A));
    int bh = (height + (TILE<<A)-1)/((TILE<<A));

    int bwOld = (width + (TILE<<(A+1))-1)/((TILE<<(A+1)));

	cudaMalloc((void**)retD,bh*bw* sizeof(int2 ));

	dim3 blocks(bw, bh);
	dim3 threads(2*SEARCHWINDOW+1,2*SEARCHWINDOW+1);

	//printf("Invocacion SERCHcu\t");
	SERCHcu<<<blocks, threads>>>(bw,bwOld,inMVD, *retD);
	//printf("Respuesta del Kernel\n");
}

//Subir a memoria del device y ejecutar algoritmo sad
void meCU(x264_t *h, int** mvX, int** mvY){
	TEXac.addressMode[0] = cudaAddressModeClamp;
	TEXan.addressMode[0] = cudaAddressModeClamp;

	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uint8_t>();
	cudaArray* currentArray, *lastArray;

	cudaMallocArray(&currentArray, 	&channelDesc, h->fenc->i_width[0], 		h->fenc->i_lines[0]);
	cudaMallocArray(&lastArray, 		&channelDesc, h->fref0[0]->i_width[0], 	h->fref0[0]->i_lines[0]);

	cudaMemcpy2DToArray(currentArray, 0, 0,
						h->fenc->plane[0], h->fenc->i_stride[0]* sizeof(uint8_t),
						h->fenc->i_width[0] * sizeof(uint8_t),  h->fenc->i_lines[0],
						cudaMemcpyHostToDevice);

	cudaMemcpy2DToArray(lastArray, 0, 0,
						h->fref0[0]->plane[0],  h->fref0[0]->i_stride[0]* sizeof(uint8_t),
						h->fref0[0]->i_width[0] * sizeof(uint8_t), h->fref0[0]->i_lines[0],
						cudaMemcpyHostToDevice);

	cudaBindTextureToArray(TEXac, 	currentArray, 	channelDesc);
	cudaBindTextureToArray(TEXan, 		lastArray, 		channelDesc);
  
	int bw2 = (h->fenc->i_width[0] + 63)/64;
	int bh2 = (h->fenc->i_lines[0] + 63)/64;

	int bw = (h->fenc->i_width[0] + 15)/16;
	int bh = (h->fenc->i_lines[0] + 15)/16;

	int2* MVFinal=NULL;

	{
		int2 *MVFinalD =NULL;

		{
			int2* mv1D = NULL;
			{
				int2 *mvsD;
				cudaMalloc((void**)&mvsD,bh2*bw2* sizeof(int2));
				cudaMemset(mvsD,0,bh2*bw2* sizeof(int2));
				EXcu(h->fenc->plane[0], h->fref0[0]->plane[0], h->fenc->i_stride[0], h->fref0[0]->i_stride[0], h->fenc->i_width[0], h->fenc->i_lines[0], &mv1D, 1,mvsD);

				cutilSafeCall(cudaFree(mvsD));
			}

			EXcu(h->fenc->plane[0], h->fref0[0]->plane[0], h->fenc->i_stride[0], h->fref0[0]->i_stride[0], h->fenc->i_width[0], h->fenc->i_lines[0], &MVFinalD, 0,mv1D);
			cutilSafeCall(cudaFree(mv1D));
		}

		MVFinal=(int2 *)malloc(bh*bw*sizeof(int2 ));
		cudaMemcpy(MVFinal, MVFinalD, bh*bw*sizeof(int2 ), cudaMemcpyDeviceToHost);

		cutilSafeCall(cudaFree(MVFinalD));
	}

	*mvX = (int*)malloc(bh*bw*sizeof(int));
	*mvY = (int*)malloc(bh*bw*sizeof(int));
	int y,x;
	for(y = 0 ; y < bh; y++){
		for(x = 0 ; x < bw ; x++){
			(*mvX)[x+y*bw] = MVFinal[x+y*bw].x;
// 			printf("%d",(*mvX)[x+y*bw]);
			(*mvY)[x+y*bw] = MVFinal[x+y*bw].y;

		}
	}
	free(MVFinal);

	cudaFreeArray(currentArray);
	cudaFreeArray(lastArray);

	cudaUnbindTexture(TEXac);
	cudaUnbindTexture(TEXan);
}


//-553910281
#ifdef __CUDACC__
}
#endif
